#include "CL/cl.h"
#include <sys/types.h>
#include <stdio.h>
#include <string.h>
#include <sys/time.h>

#include <clAmdBlas.h>

#define M (2048)
#define N (2048)
#define K (2048)

static const clAmdBlasOrder order = clAmdBlasColumnMajor;
static const clAmdBlasTranspose transA = clAmdBlasTrans;
static const clAmdBlasTranspose transB = clAmdBlasTrans;

# define timersub(a, b, result)                                         \
        do {                                                            \
                (result)->tv_sec = (a)->tv_sec - (b)->tv_sec;           \
                (result)->tv_usec = (a)->tv_usec - (b)->tv_usec;        \
                if ((result)->tv_usec < 0) {                            \
                        --(result)->tv_sec;                             \
                        (result)->tv_usec += 1000000;                   \
                }                                                       \
        } while (0)

const char *KernelSource = "\n"                                                         
"__kernel void sgemm_c(__global float *C, __global float *A, __global float *B,		\n"\
"					   __read_only image2d_t imgA, __read_only image2d_t imgB,		\n"\
"					   int m, int n, int k, int lda, int ldb, int ldc, 				\n"\
"					   int offsetA, int offsetB) 									\n"\
"{																					\n"\
"	const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP |		\n"\
"							  CLK_FILTER_NEAREST;									\n"\
"																					\n"\	
"	const  int tx = get_local_id(0);												\n"\
"	const  int ty = get_local_id(1);												\n"\
"																					\n"\
"	const int iby = get_group_id(1) * 96;											\n"\
"	const int ibx = get_group_id(0) * 96;											\n"\
"	const int idt = ty * 64 + tx;													\n"\
"																					\n"\
"	const int tx2 = idt%16;															\n"\
"	const int ty2 = idt/16;															\n"\
"																					\n"\
"	__local float Bb[16][97];														\n"\
"	__local float Abs[96][17];														\n"\
"																					\n"\
"	float xxA[6];																	\n"\
"	float xxB[6];																	\n"\
"																					\n"\
"	int trackA = offsetA + mul24( ibx + ty2, lda) + tx2;							\n"\
"	int trackB = offsetB + iby+ tx2 + mul24(ty2, ldb);								\n"\
"																					\n"\
"	int2 loc;																		\n"\
"	float4 buf;																		\n"\
"																					\n"\
"	int counter = 0;																\n"\
"																					\n"\
"	A += trackA; 																	\n"\
"	B += trackB; 																	\n"\
"																					\n"\
"	int tll = tx2; 																	\n"\
"																					\n"\
"//#pragma unroll																	\n"\
"	for(int y=0; y<6; y++){															\n"\
"		loc =(int2)(tx2, ty2 + ibx + 16*y );										\n"\
"//		loc =(int2)(ty2 + ibx + 16*y , tx2);										\n"\
"		buf = read_imagef(imgA , sampler , loc);									\n"\
"		Abs[ty2+16*y][tx2] = buf.x;													\n"\
"	}																				\n"\
"//#pragma unroll																	\n"\
"	for(int y=0; y<6; y++){															\n"\
"		loc = (int2)(iby + tx2 + 16*y, ty2);										\n"\
"//		loc = (ty2, iby + tx2 + 16*y);												\n"\
"		buf = read_imagef(imgB, sampler, loc);										\n"\
"		Bb[ty2][tx2+16*y] = buf.x;													\n"\
"	}																				\n"\
"																					\n"\
"	barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);							\n"\
"																					\n"\
"//	float *Bend = (float *) (B + k*ldb - 16*ldb);									\n"\
"	int Bend = k*ldb - 16*ldb;														\n"\
"	int loop = 0;																	\n"\
"	float Axs[6];																	\n"\
"	float Bxp[6];																	\n"\
"																					\n"\
"	float Cb[36] = {0,0,0,0,0,0, 0,0,0,0,0,0,  0,0,0,0,0,0, 0,0,0,0,0,0,			\n"\
"		0,0,0,0,0,0, 0,0,0,0,0,0};													\n"\
"	do 																				\n"\
"	{																				\n"\
"		tll+=16;																	\n"\
"		A += 16;																	\n"\
"		B += 16*ldb;																\n"\
"		counter += 16;																\n"\
"		loop += 16*ldb;																\n"\
"//#pragma unroll																	\n"\
"		for( int y=0; y<6; y++){													\n"\
"			loc =(int2)(tx2 + counter , ty2 + ibx + 16*y);							\n"\
"//			loc = (int2)(ty2 + ibx + 16*y, tx2 + counter);							\n"\
"			buf = read_imagef(imgA, sampler, loc);									\n"\
"			xxA[y] = buf.x;															\n"\
"		}																			\n"\
"//#pragma unroll																	\n"\
"		for( int y=0; y<6; y++){													\n"\
"			loc = (int2)(iby + tx2 + 16*y, ty2 + counter);							\n"\
"//			loc = (int2)(ty2 + counter, iby + tx2 + 16*y);							\n"\
"			buf = read_imagef(imgB, sampler, loc);									\n"\
"			xxB[y] = buf.x;															\n"\
"		}																			\n"\
"//#pragma unroll 																	\n"\
"		for( int j1=0;j1<16;j1++)													\n"\
"		{																			\n"\
"//#pragma unroll																	\n"\
"			for( int y=0; y<6; y++)													\n"\
"				Axs[y] =  Abs[tx2 + y*16][j1];										\n"\
"																					\n"\
"//#pragma unroll																	\n"\
"			for( int y=0; y<6; y++)													\n"\
"				Bxp[y]= Bb[j1][ty2 + y*16];											\n"\
"																					\n"\
"//#pragma unroll 																	\n"\
"			for( int x=0; x<6; x++)													\n"\
"//#pragma unroll 																	\n"\
"				for( int y=0;y<6;y++)												\n"\
"					Cb[x*6+y] += Axs[x]*Bxp[y];										\n"\
"		}																			\n"\
"																					\n"\
"		barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);						\n"\
"																					\n"\
"//#pragma unroll																	\n"\
"		for( int y=0; y<6; y++)														\n"\
"			Abs[ty2 + 16*y][tx2] = xxA[y];											\n"\
"																					\n"\
"//#pragma unroll																	\n"\
"		for( int y=0; y<6; y++)														\n"\
"			Bb[ty2][tx2+y*16] = xxB[y];												\n"\
"																					\n"\
"		barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);						\n"\
"																					\n"\
"	} 																				\n"\
"	while (loop < Bend);															\n"\
"																					\n"\
"	C += tx2 + ibx  + mul24 (ty2 +  iby ,ldc);										\n"\
"																					\n"\
"//#pragma unroll 																	\n"\
"	for( int j1=0; j1<16; j1++)														\n"\
"	{																				\n"\
"//#pragma unroll																	\n"\
"		for( int y=0; y<6; y++)														\n"\
"			Axs[y] =  Abs[tx2 + y*16][j1];											\n"\
"																					\n"\
"//#pragma unroll																	\n"\
"		for( int y=0; y<6; y++)														\n"\
"			Bxp[y]= Bb[j1][ty2 + y*16];												\n"\
"																					\n"\
"//#pragma unroll 																	\n"\
"		for( int x=0; x<6; x++)														\n"\
"//#pragma unroll 																	\n"\
"			for( int y=0; y<6; y++)													\n"\
"				Cb[x*6+y]  += Axs[x]*Bxp[y];										\n"\
"	}																				\n"\
"																					\n"\
"	int gy = iby + ty2;																\n"\
"//#pragma unroll																	\n"\
"	for( int y=0; y<6; y++, gy+=16)													\n"\
"	{																				\n"\
"		int gx = ibx + tx2; 														\n"\
"//#pragma unroll																	\n"\
"		for(int x=0; x<6; x++, gx+=16)												\n"\
"		{																			\n"\
"			if (gx < m && gy < n)													\n"\
"				C[x*16] = Cb[y+x*6] + 0 * C[x*16];									\n"\
"		}																			\n"\
"																					\n"\
"		C+=ldc*16;																	\n"\
"	}																				\n"\
"																					\n"\
"}																					\n"\
"\n";

size_t shrRoundUp(int group_size, int global_size)
{
	int r = global_size % group_size;
	if(r == 0){
		return global_size;
	}
	else
	{
		return global_size + group_size - r;
	}
}

int main()
{
	struct timeval time_start, time_end, time_res;

	int n = N;
	int m = M;
	int k = K;

	int mk = M*K;
	int kn = K*N;
	int mn = M*N;

	int lda = K;
	int ldb = N;
	int ldc = M;

	int offsetA = 0;
	int offsetB = 0;

	int i;
	int j;

	int count = 0;
	
	float alpha = 1;
	float beta = 0;

	float *a,*b,*c, *res;

	cl_int err;
	
	int size = 2048;
	size_t local[2] = {64, 4};
	size_t global[2] = {1408,88};

	cl_platform_id platform;
	cl_device_id device;
	cl_context context;
	cl_command_queue command;
	cl_program program;
	cl_kernel kernel;

	cl_event event = NULL;

	size_t origin[3];
	size_t region[3];

	cl_image_format format = {CL_R, CL_FLOAT};	
	cl_image_desc desc;

	cl_mem img_a;
	cl_mem img_b;
	cl_mem dev_a;
	cl_mem dev_b;
	cl_mem dev_c;
	cl_mem dev_res;

	a = (float*)malloc(mk * sizeof(float));
	b = (float*)malloc(kn * sizeof(float));
	c = (float*)malloc(mn * sizeof(float));
	res = (float*)malloc(mn * sizeof(float));

	origin[0] = 0;
	origin[1] = 0;
	origin[2] = 0;
	region[0] = N;
	region[1] = M;
	region[2] = 1;


        for (i = 0 ; i < mk ; i++){
//		a[i] = i;
	       	a[i] = rand() / (float)RAND_MAX;
    	}

    	for (i = 0 ; i < kn ; i++){
//		b[i] = kn-i;
        	b[i] = rand() / (float)RAND_MAX;
    	}
/*
    	for (i = 0 ; i < mn ; i++){
//		c[i] = 1;
        	c[i] = rand() / (float)RAND_MAX;
			res[i] = c[i];
    	}
*/
	size_t row_pitch = sizeof(float) * N;
	size_t slice_pitch = 0;

	//get platform id
	err = clGetPlatformIDs(1, &platform, NULL);
	if(err!= CL_SUCCESS){
		printf( "clGetPlatformIDs() failed with %d\n", err );
		exit(-1);
	}

	//get device id
	err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
	if(err!= CL_SUCCESS){
		printf( "clGetDeviceIDs() failed with %d\n", err );
		exit(-1);
	}

	char buf[256];

	clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION , sizeof(char) * 256, buf, NULL);
	printf("%s\n",buf);

	//Create Context
	context = clCreateContext(0, 1, &device, NULL, NULL, &err);
	if(err != CL_SUCCESS){
		printf( "clCreateContext() failed with %d\n", err );
		exit(-1);
	}

	//create command queue
	command = clCreateCommandQueue(context, device, 0, &err);
	if(err != CL_SUCCESS){
		printf( "clCreateCommandQueue() failed with %d\n", err );
		exit(-1);
	}
	
	// Setup clAmdBlas
	err = clAmdBlasSetup();
	if(err != CL_SUCCESS){
		printf("clAmdBlasSetup() failed with %d\n",err);
		clReleaseCommandQueue(command);
		clReleaseContext(context);
		exit(-1);
	}	


	//Create program
	program = clCreateProgramWithSource(context, 1, (const char **)& KernelSource, NULL, &err);
	if(err != CL_SUCCESS){
		printf( "clCreateProgramWithSource() failed with %d\n", err );
		exit(-1);
	}

	//build program
	err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (err != CL_SUCCESS) 
	{
			printf( "clBuildProgram() failed with %d\n", err );
			if(err == CL_BUILD_PROGRAM_FAILURE){
					size_t log_size;
					char *log;
					clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
					log = (char *)malloc(log_size);
					clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
					printf("%s\n", log);
					exit(-1);
			}
			exit(-1);
	}


	//create kernel
	kernel = clCreateKernel(program, "sgemm_c", &err);
	if (err != CL_SUCCESS) {
		printf( "clBuildKernel() failed with %d\n", err );
                exit(-1);
        }
	
	//create buffer
	img_a = clCreateImage2D(context, CL_MEM_READ_ONLY, &format, K, M, 0 , NULL, &err);
	if (err != CL_SUCCESS) {
		printf("Can't Create image a! %d\n",err);
		exit(-1);
	}
	img_b = clCreateImage2D(context, CL_MEM_READ_ONLY, &format, N, K, 0, NULL, &err);
	if (err != CL_SUCCESS) {
		printf("Can't Create image b! %d\n",err);
		exit(-1);
	}
	dev_a = clCreateBuffer(context, CL_MEM_READ_WRITE, M * K * sizeof(float),NULL,&err);
	if (err != CL_SUCCESS) {
		printf("Can't Create buffer a! %d\n",err);
		exit(-1);
	}
	dev_b = clCreateBuffer(context, CL_MEM_READ_WRITE, K * N * sizeof(float),NULL,&err);
	if (err != CL_SUCCESS) {
		printf("Can't Create buffer b! %d\n",err);
		exit(-1);
	}
	dev_c = clCreateBuffer(context, CL_MEM_READ_WRITE, M * N * sizeof(float),NULL,&err);
	if (err != CL_SUCCESS) {
		printf("Can't Create buffer c! %d\n",err);
		exit(-1);
	}
	
	dev_res = clCreateBuffer(context, CL_MEM_READ_WRITE, M * N * sizeof(float),NULL,&err);
	if (err != CL_SUCCESS) {
		printf("Can't Create buffer c! %d\n",err);
		exit(-1);
	}
	//enqueue wrte buffer
	
	err = clEnqueueWriteBuffer(command, dev_a, CL_TRUE, 0, M * K * sizeof(float),a,0 ,NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Write Buffer A! %d\n",err);
		exit(-1);
	}
	err = clEnqueueWriteBuffer(command, dev_b, CL_TRUE, 0, K * N * sizeof(float),b,0 ,NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Write Buffer B! %d\n",err);
		exit(-1);
	}
	
	err = clEnqueueWriteImage(command, img_a, CL_TRUE, origin, region, row_pitch, slice_pitch, a, 0, NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Write image a ! %d\n",err);
		exit(-1);
	}
	err = clEnqueueWriteImage(command, img_b, CL_TRUE, origin, region, row_pitch, slice_pitch, b, 0, NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Write image b ! %d\n",err);
		exit(-1);
	}
	err = clEnqueueWriteBuffer(command, dev_c, CL_TRUE, 0, M * N * sizeof(float),c,0 ,NULL, NULL);
	if (err != CL_SUCCESS) {
		printf("Can't Enqueue Write Buffer C! %d\n",err);
		exit(-1);
	}
	
	
//set function arguments

	err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &dev_c);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment c!\n");
		exit(-1);
	}
	err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &dev_a);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment a!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &dev_b);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment b!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 3, sizeof(cl_mem), &img_a);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment a!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 4, sizeof(cl_mem), &img_b);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment b!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 5, sizeof(int), &m);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment m!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 6, sizeof(int), &n);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment n!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 7, sizeof(int), &k);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment k!\n");
		exit(-1);
	}
	err = clSetKernelArg(kernel, 8, sizeof(int), &lda);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment lda!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 9, sizeof(int), &ldb);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment ldb!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 10, sizeof(int), &ldc);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment ldc!\n");
		exit(-1);
	}

	err = clSetKernelArg(kernel, 11, sizeof(int), &offsetA);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment offsetA!\n");
		exit(-1);
	}
	err = clSetKernelArg(kernel, 12, sizeof(int), &offsetB);
	if (err != CL_SUCCESS) {
		printf("Can't Set Kernel Argment offsetB!\n");
		exit(-1);
	}
	gettimeofday(&time_start, NULL);
	
	err = clEnqueueNDRangeKernel(command, kernel, 2, NULL,global,local, 0, NULL, &event);

//	err = clWaitForEvents(1, &event);

	if (err != CL_SUCCESS) {
                printf("Can't Enqueue NDRange Kernel %d!\n",err);
                exit(-1);
        }


	err = clFinish(command);
	if (err != CL_SUCCESS) {
                printf("Can't finish!\n");
                exit(-1);
        }

	gettimeofday(&time_end,NULL);

	timersub(&time_end, &time_start, &time_res);

	printf("%ld,%03ld ms\n", time_res.tv_sec, time_res.tv_usec/1000);
	
	err = clEnqueueReadBuffer( command, dev_c, CL_TRUE, 0,N * N * sizeof(float) , c , 0, NULL, NULL);
	if (err != CL_SUCCESS) {
                printf("Can't Enqueue Read Buffer\n");
                exit(-1);
        }

	err = clAmdBlasSgemm(order, transA, transB, M, N, K, alpha, dev_a, K, dev_b,
				N, beta, dev_res, N, 1, &command, 0, NULL, &event);
	if(err != CL_SUCCESS){
		printf("clAmdBlasSgemm() failed with %d\n", err);
		exit(-1);
	}
	else{
		err = clWaitForEvents(1, &event);
		err = clEnqueueReadBuffer(command, dev_res, CL_TRUE, 0,
					M * N * sizeof(*res),
					res, 0, NULL, NULL);
	}

//	printf("%lf %lf %lf %lf %lf \n",c[0],c[1],c[2],c[3],c[4]);
//	printf("%lf %lf %lf %lf %lf \n",res[0],res[1],res[2],res[3],res[4]);

	for(i = 0; i < N ; i++){
		for(j = 0 ; j < N ; j++){
			if( (c[i*N+j] - res[i*N+j]) > 0.001 || (c[i*N+j] - res[i*N+j]) < -0.001){
                	count++;
            		}
        	}
    	}

    	printf("number of differences : %d\n",count);

	
	clReleaseMemObject(dev_a);
	clReleaseMemObject(dev_b);
	clReleaseMemObject(dev_c);
	clReleaseMemObject(dev_res);

	clReleaseCommandQueue(command);
	clReleaseContext(context);

    return 0;
}
